/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/cuqu.h>

#include <cuos/config.h>
#include <cuos/types.h>
#include <cuos/detail/system_services.h>

namespace cuos
{
    typedef cuqu::host::queue<sys_req_t>   sys_req_queue_t;
    typedef sys_req_queue_t::event_t       sys_req_event_t;

    typedef cuqu::host::queue<sys_rep_t>   sys_rep_queue_t;
    typedef sys_rep_queue_t::event_t       sys_rep_event_t;

    namespace device {

        //const int tmout = cuqu::timeout_infinite;
        enum constants {
            //tmout_ms = 1000
            tmout_ms = cuqu::timeout_infinite
        };

        struct sys_queue {
            cuqu::device::queue *m_req;
            cuqu::device::queue *m_rep;
            bool                 m_initialized;

            __device__ void dump_queues() {
                cuqu::detail::__queue_dump_details(m_req);
                cuqu::detail::__queue_dump_details(m_rep);
            }

            // blocking execution
            __device__ int exec(sys_req_event_t *gbl_req, sys_rep_event_t *gbl_rep) {
                const unsigned int tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
                const unsigned int bid = blockIdx.x + gridDim.x*blockIdx.y;
                const unsigned int block_size = blockDim.x*blockDim.y*blockDim.z;
                //const unsigned int grid_size = gridDim.y*gridDim.x;
                const unsigned int gid = tid + bid*block_size;

                int retcode = cuqu::SUCCESS;
                //cuqu_error("this=%p gbl_req=%p m_req=%p\n", this, gbl_req, m_req);
#if CUOS_ENABLE_EXTRA_CHECKS
                if(!m_initialized || !m_req || !m_rep) {
                    if(0 == gid) {
                        cuqu_error("ERROR invalid sys_queue %p state m_initialized=%d m_req=%p m_rep=%p\n", this, m_initialized, m_req, m_rep);
                    }
                    return RUNTIME;
                }
#endif
#if CUOS_ENABLE_EXTRA_CHECKS
                unsigned int reqerrs = cuqu::detail::__queue_n_failures(m_req);
                unsigned int reperrs = cuqu::detail::__queue_n_failures(m_rep);
                if(reqerrs > 0 || reperrs > 0) {
                    cuqu_error("[%d:%d] ERROR req:%d rep:%d errors in queue\n", bid, tid, 
                               reqerrs, reperrs );
                    return RUNTIME;                    
                }
#endif

                retcode = m_req->timed_push(gbl_req, tmout_ms);

                if(retcode != cuqu::SUCCESS) {
                    cuqu_error("ERROR %d in req->timed_push()", retcode);
                    return retcode;
                }

                retcode = m_rep->timed_fetch(gbl_rep, tmout_ms);

                if(retcode != cuqu::SUCCESS) {
                    cuqu_error("ERROR %d in rep->timed_fetch()", retcode);
                }

#if CUOS_ENABLE_EXTRA_CHECKS
                using namespace cuos::system_service;
                sys_rep_t rep = gbl_rep->u.t;
                if(0 == gid && (rep.serial != 0 || rep.retval.tag != param::tag_int)) {
                    cuos_error("[%d:%d] error in reply format: serial=%x flags=%x tag=%x count=%x data=%x\n",
                               bid, tid, rep.serial, rep.flags, rep.retval.tag, rep.retval.count, rep.retval.si);
                    dump_queues();
                }
#endif

                return retcode;
            }

            __device__ int exec_oneway(sys_req_event_t *gbl_req) {
                const unsigned int tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
                const unsigned int bid = blockIdx.x + gridDim.x*blockIdx.y;
                const unsigned int block_size = blockDim.x*blockDim.y*blockDim.z;
                //const unsigned int grid_size = gridDim.y*gridDim.x;
                const unsigned int gid = tid + bid*block_size;

                int retcode = cuqu::SUCCESS;
                //cuqu_error("this=%p gbl_req=%p m_req=%p\n", this, gbl_req, m_req);
                if(!m_initialized || !m_req || !m_rep) {
                    if(0 == gid) {
                        cuqu_error("ERROR invalid sys_queue %p state m_initialized=%d m_req=%p m_rep=%p\n", this, m_initialized, m_req, m_rep);
                    }
                    return RUNTIME;
                }
                retcode = m_req->timed_push(gbl_req, tmout_ms);
                return retcode;
            }
#if CUOS_NOT_IMPLEMENTED_YET
            // split execution, push + wait
            __device__ int push_async(gpu_event_t *gbl_desc) {
                return SUCCESS;
            }
            __device__ int wait_end(gpu_event_t *gbl_desc) {
                return SUCCESS;
            }
            __device__ int test_end(gpu_event_t *gbl_desc) {
                return SUCCESS;
            }
#endif
        };
        typedef struct sys_queue sys_queue_t;

        //__device__ req_queue_t req_queue();
        //__constant__ __device__  sys_queue_t sys_q;
        //extern __device__  sys_queue_t *ptr_sys_q;
        //__device__ sys_queue_t *sys_q();

        __device__ static inline int syscall_exec(sys_queue_t *sys_q, sys_req_event_t *gbl_req, sys_rep_event_t *gbl_rep) {
            //const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
            //const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
            //const uint block_size = blockDim.x*blockDim.y*blockDim.z;
            //const uint grid_size = gridDim.y*gridDim.x;
            //const uint gid = tid + bid*block_size;
            //const uint tot_threads = block_size*grid_size;
            return sys_q->exec(gbl_req, gbl_rep);
        }

        __device__ static inline int syscall_exec_oneway(sys_queue_t *sys_q, sys_req_event_t *gbl_req) {
            //const uint tid = threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
            //const uint bid = blockIdx.x + gridDim.x*blockIdx.y;
            //const uint block_size = blockDim.x*blockDim.y*blockDim.z;
            //const uint grid_size = gridDim.y*gridDim.x;
            //const uint gid = tid + bid*block_size;
            //const uint tot_threads = block_size*grid_size;
            return sys_q->exec_oneway(gbl_req);
        }

    } // end namespace device

    typedef device::sys_queue_t sys_queue_t;

} // end namespace cuos

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
